/*
 *  Copyright 2009 Advanced Industrial Science and Technology (AIST).
 *
 *  Licensed under the Apache License, Version 2.0 (the "License");
 *  you may not use this file except in compliance with the License.
 *  You may obtain a copy of the License at
 *
 *      http://www.apache.org/licenses/LICENSE-2.0
 *
 *  Unless required by applicable law or agreed to in writing, software
 *  distributed under the License is distributed on an "AS IS" BASIS,
 *  WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
 *  See the License for the specific language governing permissions and
 *  limitations under the License.
 */

/*! \file
 *  \brief
 */
#pragma once

#include <thrust/iterator/iterator_traits.h>
#include <thrust/detail/type_traits.h>

namespace thrust
{
namespace detail
{
namespace device
{
namespace cuda
{
namespace detail
{

/*
 * Extend a n-elements vector to a N-elements vector ( N > n) by keyvalue
 *
 * 
 */

template <typename InputIterator1,
          typename InputIterator2,
          typename OutputIterator>
__global__ void
reduce_by_keyvalue( InputIterator1 key,
                    InputIterator2 value,
                    const unsigned int n,
                    OutputIterator *result,
                    const unsigned int N // N is size of result elements( =num_vertices)
		    )
{
  const unsigned int grid_size =blockDim.x * gridDim.x;
  unsigned int tid =blockDim.x * blockIdx.x + threadIdx.x;

  key   +=tid;
  value +=tid;
  
  for( ; tid < n; tid +=grid_size){
  
    unsigned int m_key =thrust::detail::device::dereference(key);
    unsigned int m_value =thrust::detail::device::dereference(value);
    if( m_key <N)
    result[m_key] =m_value;
    
    key +=grid_size;
    value +=grid_size;
  }
}

} // end namespace detail

template <typename InputIterator1,
          typename InputIterator2,
          typename OutputIterator>
void
reduce_by_keyvalue( InputIterator1 key_first,
                    InputIterator1 key_last,
                    InputIterator2 value_first,
                    OutputIterator result_first,
                    OutputIterator result_last)
{
  
  typedef typename thrust::iterator_traits<InputIterator1>::difference_type difference_type;
  
  difference_type n =key_last - key_first;
  difference_type N =result_last - result_first;
  
  const std::size_t BLOCK_SIZE =512; // number of threads per block
  const std::size_t MAX_BLOCKS =thrust::detail::device::cuda::arch::max_active_threads() / BLOCK_SIZE;
  const std::size_t NUM_BLOCKS =std::min( MAX_BLOCKS, n + (BLOCK_SIZE - 1) / BLOCK_SIZE);

  typedef typename thrust::iterator_traits<OutputIterator>::value_type OutputType;
  thrust::detail::raw_cuda_device_buffer<OutputType> temp( result_first, result_last);
  
  detail::reduce_by_keyvalue<<<NUM_BLOCKS, BLOCK_SIZE>>>( key_first, 
                                                          value_first, 
                                                          n, 
                                                          raw_pointer_cast(&temp[0]), 
                                                          N);
  thrust::copy( temp.begin(), temp.end(), result_first);
  
}

} // end namespace cuda
} // end namespace device
} // end namespace detail
} // end namespace thrust
